
#define HASH0_(from, to, modulo) \
(to) = (from);           \
(to) -= ((to)<<6);    \
(to) ^= ((to)>>17);   \
(to) -= ((to)<<9);    \
(to) ^= ((to)<<4);    \
(to) -= ((to)<<3);    \
(to) ^= ((to)<<10);   \
(to) ^= ((to)>>15);   \
(to) %= (modulo);     \
(to) <<= 1;

#define HASH1_(from, to, modulo)             \
(to) = (from);                              \
(to) ^= ((to) >> 20) ^ ((to) >> 12);        \
(to) = (to) ^ ((to) >> 7) ^ ((to) >> 4);    \
(to) %= (modulo);                           \
(to) <<= 1;

__kernel void cockooFind(__global uint * v, __global uint * out,__global uint * hashTable, __local uint * founded, uint size,__global uint * gdbg ){

        int gid = get_global_id(0);
        int lid = get_local_id(0);
        __global uint * res = out + get_group_id(0) * get_local_size(0);

        uint i = *(v + gid);

        uint pos;
        uint targetIdx;// = -1;
        uint NO_TARGET = -1;

        if (lid == 0){
            *founded = 0;
        }
        barrier( CLK_LOCAL_MEM_FENCE );


        HASH0_( i, pos, size );
        targetIdx = i == *(hashTable + pos) ? atomic_inc( founded ) : NO_TARGET;
        HASH1_( i, pos, size );
        pos += 1;
        targetIdx = i == *(hashTable + pos) ? atomic_inc( founded ) : targetIdx;

        barrier( CLK_LOCAL_MEM_FENCE );

        if ( targetIdx != NO_TARGET ){
            //save value that was counted by this thread
            *(res + targetIdx) = i;
        }

        if (lid == 0){
            //save sum
            gdbg[get_group_id(0)] = (*founded);
        }
        if( get_global_id(0) == 0 )
        {
         gdbg[get_num_groups(0) + 0] = get_num_groups(0);
         gdbg[get_num_groups(0) + 1] = get_global_size(0);
         gdbg[get_num_groups(0) + 2] = get_local_size(0);
        }

}
